Si uno desea que el kernel imprima algo mientras se ejecuta se tiene que hacer uso de printf que es la función utilizada en C y CUDA para imprimir. Sin embargo hay recordar algunas cosas. En primer lugar, printf imprime cosas en la consola, es por eso que el lector notará al ejecutar la siguiente celda, que incluso cuando en el kernel le estamos ordenando que imprima, en el notebook no obtenemos nada. A primera vista parecería un error garrafal, sin embargo si comentáramos la primer linea (la que comienza con %%writefile), al revisar la consola desde la cuál abrimos el notebook se observaría que los resultados esperados se han impreso ahí. Ejecutemos la siguiente celda, con el comando mágico %%writefile para escribir el contenido de la celda en un archivo.
In [1]:
%%writefile ./Programas/saludar.py
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule
mod = SourceModule("""
#include <stdio.h>
__global__ void saluda()
{
printf("Mi indice x es %d, mi indice en y es %d\\n", threadIdx.x, threadIdx.y);
}
""")
func = mod.get_function("saluda")
func(block=(4,4,1))
Ahora ejecutemos el programa
In [2]:
!ipython ./Programas/saludar.py
Por último hagamos un kernel que imprima más datos acerca del thread, como su índice de bloque, la dimensión del bloque, y sus índices dentro del bloque.
In [3]:
%%writefile ./Programas/saludar_bloques.py
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule
mod = SourceModule("""
#include <stdio.h>
__global__ void say_hi()
{
printf("Soy el thread numero %d en threadIdx.x:%d.threadIdx.y:%d blockIdx.:%d blockIdx.y:%d blockDim.x:%d blockDim.y:%d\\n",(threadIdx.x+threadIdx.y*blockDim.x+(blockIdx.x*blockDim.x*blockDim.y)+(blockIdx.y*blockDim.x*blockDim.y)),threadIdx.x, threadIdx.y,blockIdx.x,blockIdx.y,blockDim.x,blockDim.y);
}
""")
func = mod.get_function("say_hi")
func(block=(4,4,1),grid=(2,2,1))
Lo ejecutamos. Esperamos que se ejecuten uno por uno los bloques; veamos qué es lo que pasa.
In [4]:
!ipython ./Programas/saludar_bloques.py
En efecto se ejecutaron uno a la vez.
In [5]:
import pycuda.driver as drv
import pycuda.tools
import pycuda.autoinit
import numpy
from pycuda.compiler import SourceModule
import pycuda.gpuarray as gpuarray
import pycuda.cumath
from pycuda.elementwise import ElementwiseKernel
blocks = 64
block_size = 128
valores = blocks * block_size
print "Usando", valores, "valores"
# Número de iteraciones para los cálculos
n_iter = 100000
print "Calculando %d iteraciones" % (n_iter)
# Crear dos timers
inicio = drv.Event()
fin = drv.Event()
# SourceModele
mod = SourceModule("""
__global__ void gpusin(float *dest, float *a, int n_iter)
{
const int i = blockDim.x*blockIdx.x + threadIdx.x;
for(int n = 0; n < n_iter; n++) {
a[i] = sin(a[i]);
}
dest[i] = a[i];
}
""")
gpusin = mod.get_function("gpusin")
# creamos un arreglo 1s
a = numpy.ones(valores).astype(numpy.float32)
# creamos un arreglo para guardar el resultado
dest = numpy.zeros_like(a)
inicio.record() # comenzamos a tomar el tiempo
gpusin(drv.Out(dest), drv.In(a), numpy.int32(n_iter), grid=(blocks,1), block=(block_size,1,1) )
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con SourceModule y primeros 3 resultados:"
print "%fs, %s" % (segs, str(dest[:3]))
# Usando ElementwiseKernel con sin en un ciclo for en C
kernel = ElementwiseKernel(
"float *a, int n_iter",
"for(int n = 0; n < n_iter; n++) { a[i] = sin(a[i]);}",
"gpusin")
a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
kernel(a_gpu, numpy.int(n_iter))
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con Elementwise y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))
# Elementwise haciendo el loop en Python
kernel = ElementwiseKernel(
"float *a",
"a[i] = sin(a[i]);",
"gpusin")
a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
for i in range(n_iter):
kernel(a_gpu)
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con Elementwise en Python y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))
# GPUArray
# El resultado se copia a la memoria principal en cada iteración (esto es un cuello de botella)
a = numpy.ones(valores).astype(numpy.float32)
a_gpu = gpuarray.to_gpu(a)
inicio.record() # comenzamos a tomar el tiempo
for i in range(n_iter):
a_gpu = pycuda.cumath.sin(a_gpu)
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con GPUArray y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a_gpu.get()[:3]))
# CPU
a = numpy.ones(valores).astype(numpy.float32)
inicio.record() # comenzamos a tomar el tiempo
inicio.synchronize()
for i in range(n_iter):
a = numpy.sin(a)
fin.record() # terminamos de tomar el tiempo
# calculamos cuánto duró la corrida
fin.synchronize()
segs = inicio.time_till(fin)*1e-3
print "Tiempo con CPU y primeros 3 resultados:"
print "%fs, %s" % (segs, str(a[:3]))
Como podemos ver, no siempre es más eficiente hacer las cosas en el GPU; obviamente esta vez no lo fue porque forzamos al código a ser ineficiente, sin embargo es usual que el programador introduzca cuellos de botella como el que introdujimos sin siquiera darse cuenta.
En los materiales adicionales se agregan unas lecturas que pueden resultar ilustrativas de este punto.